extern "C"
__global__ void fft(int size, int N, float* a){
	__shared__ float wlist[WLIST_SIZE*2];
		
	float w1[4];
	float w2[4];
	float tmp[4];
	float tmp2[4];
	int step = 0;
	int step2 = 0;
	int flag = 0;
	int wiInd = 0;
	int wjInd = 0;
	int dsInd = 0;
	int rInd = 0;
	int rsInd = 0;
	int rMirr = 0;
	int dMirr = 0;
	int rightIndex = 0;
	int downIndex = 0;
	
	int idx = threadIdx.x/4;
	int blockX = blockDim.x/4;
		
	float a00a;
	float a00b;
	float a00c;
	float a00d;
	float a01a;
	float a01b;
	float a01c;
	float a01d;
	float a10a;
	float a10b;
	float a10c;
	float a10d;
	float a11a;
	float a11b;
	float a11c;
	float a11d;
	float power;
		
	int a00ai;
	int a00bi;
	int a00ci;
	int a00di;
	int a01ai;
	int a01bi;
	int a01ci;
	int a01di;
	int a10ai;
	int a10bi;
	int a10ci;
	int a10di;
	int a11ai;
	int a11bi;
	int a11ci;
	int a11di;
	
	
	if (threadIdx.y == 0)
	{
		for (int k = threadIdx.x; k < N; k = k + blockDim.x) {
			power = -2*3.14159265358979323846*k/N;
			wlist[2*k] = cos(power);
			wlist[2*k+1] = sin(power);
		}
	}
	__syncthreads();
		
	for (int q = 0; q < size; q++) {
		step = ceil(pow(2.0f,q));
		step2 = step * 2;
		for (int down = threadIdx.y; down < N - step; down = down + blockDim.y) {
			for (int right = idx; right < N - step; right = right + blockX) {
				//flag = (down/step) % 2;
				//flag = flag + (right/step) % 2;
				flag = ((down % step2) <= (step2/4)) ? 0 : 1;
				flag = ((right % step2) <= (step2/4)) ? flag : 1;
				if (flag == 0){
					if (threadIdx.x % 4 == 0){
						wiInd = (down % step2)*(N/step2) * 2;
						wjInd = (right % step2)*(N/step2) * 2;
						
						w1[0] = wlist[wiInd];
						w1[1] = wlist[wiInd+1];
						w2[0] = wlist[wjInd];
						w2[2] = wlist[wjInd+1];
							
						dsInd = down + step;
						rInd = 4*right;
						rsInd = 4*(right+step);
							
						a00ai = down*N*4 + rInd;
						a00bi = a00ai + 1;
						a00ci = a00ai + 2;
						a00di = a00ai + 3;
						a01ai = down*N*4 + rsInd;
						a01bi = a01ai + 1;
						a01ci = a01ai + 2;
						a01di = a01ai + 3;
						a10ai = dsInd*N*4 + rInd;
						a10bi = a10ai + 1;
						a10ci = a10ai + 2;
						a10di = a10ai + 3;
						a11ai = dsInd*N*4 + rsInd;
						a11bi = a11ai + 1;
						a11ci = a11ai + 2;
						a11di = a11ai + 3;
						
						//XXX = XXX.add(a00);
						a00a = a[a00ai];
						a00b = a[a00bi];
						a00c = a[a00ci];
						a00d = a[a00di];
						a10a = a[a00ai];
						a10b = a[a00bi];
						a10c = a[a00ci];
						a10d = a[a00di];
						a01a = a[a00ai];
						a01b = a[a00bi];
						a01c = a[a00ci];
						a01d = a[a00di];
						a11a = a[a00ai];
						a11b = a[a00bi];
						a11c = a[a00ci];
						a11d = a[a00di];
							
						//XXX = XXX.add(t1.multy(a10));
						tmp[0] = w1[0]*a[a10ai] - w1[1]*a[a10bi];
						a00a = a00a + tmp[0];
						a10a = a10a - tmp[0];
						a01a = a01a + tmp[0];
						a11a = a11a - tmp[0];
						tmp[1] = w1[0]*a[a10bi] + w1[1]*a[a10ai];
						a00b = a00b + tmp[1];
						a10b = a10b - tmp[1];
						a01b = a01b + tmp[1];
						a11b = a11b - tmp[1];
						tmp[2] = w1[0]*a[a10ci] - w1[1]*a[a10di];
						a00c = a00c + tmp[2];
						a10c = a10c - tmp[2];
						a01c = a01c + tmp[2];
						a11c = a11c - tmp[2];
						tmp[3] = w1[0]*a[a10di] + w1[1]*a[a10ci];							
						a00d = a00d + tmp[3];
						a10d = a10d - tmp[3];
						a01d = a01d + tmp[3];
						a11d = a11d - tmp[3];
							
						//XXX = XXX.add(t2.multy(a01));
						tmp[0] = w2[0]*a[a01ai] - w2[2]*a[a01ci];
						a00a = a00a + tmp[0];
						a10a = a10a + tmp[0];
						a01a = a01a - tmp[0];
						a11a = a11a - tmp[0];
						tmp[1] = w2[0]*a[a01bi] - w2[2]*a[a01di];
						a00b = a00b + tmp[1];
						a10b = a10b + tmp[1];
						a01b = a01b - tmp[1];
						a11b = a11b - tmp[1];
						tmp[2] = w2[0]*a[a01ci] + w2[2]*a[a01ai];
						a00c = a00c + tmp[2];
						a10c = a10c + tmp[2];
						a01c = a01c - tmp[2];
						a11c = a11c - tmp[2];
						tmp[3] = w2[0]*a[a01di] + w2[2]*a[a01bi];	
						a00d = a00d + tmp[3];
						a10d = a10d + tmp[3];
						a01d = a01d - tmp[3];
						a11d = a11d - tmp[3];
						
						//XXX = XXX.add(t1.multy(t2.multy(a11)));
						tmp[0] = w2[0]*a[a11ai] - w2[1]*a[a11bi] - w2[2]*a[a11ci] + w2[3]*a[a11di];
						tmp[1] = w2[0]*a[a11bi] + w2[1]*a[a11ai] - w2[2]*a[a11di] + w2[3]*a[a11ci];
						tmp[2] = w2[0]*a[a11ci] - w2[1]*a[a11di] + w2[2]*a[a11ai] - w2[3]*a[a11bi];
						tmp[3] = w2[0]*a[a11di] + w2[1]*a[a11ci] + w2[2]*a[a11bi] + w2[3]*a[a11ai];	
						tmp2[0] = w1[0]*tmp[0] - w1[1]*tmp[1] - w1[2]*tmp[2] + w1[3]*tmp[3];
						tmp2[1] = w1[0]*tmp[1] + w1[1]*tmp[0] - w1[2]*tmp[3] + w1[3]*tmp[2];
						tmp2[2] = w1[0]*tmp[2] - w1[1]*tmp[3] + w1[2]*tmp[0] - w1[3]*tmp[1];
						tmp2[3] = w1[0]*tmp[3] + w1[1]*tmp[2] + w1[2]*tmp[1] + w1[3]*tmp[0];	
						a00a = a00a + tmp2[0];
						a00b = a00b + tmp2[1];
						a00c = a00c + tmp2[2];
						a00d = a00d + tmp2[3];
						a10a = a10a - tmp2[0];
						a10b = a10b - tmp2[1];
						a10c = a10c - tmp2[2];
						a10d = a10d - tmp2[3];
						a01a = a01a - tmp2[0];
						a01b = a01b - tmp2[1];
						a01c = a01c - tmp2[2];
						a01d = a01d - tmp2[3];
						a11a = a11a + tmp2[0];
						a11b = a11b + tmp2[1];
						a11c = a11c + tmp2[2];
						a11d = a11d + tmp2[3];
							
						a[a00ai] = a00a;
						a[a00bi] = a00b;
						a[a00ci] = a00c;
						a[a00di] = a00d;
						a[a10ai] = a10a;
						a[a10bi] = a10b;
						a[a10ci] = a10c;
						a[a10di] = a10d;
						a[a01ai] = a01a;
						a[a01bi] = a01b;
						a[a01ci] = a01c;
						a[a01di] = a01d;
						a[a11ai] = a11a;
						a[a11bi] = a11b;
						a[a11ci] = a11c;
						a[a11di] = a11d;
					}	
					
					//MIRROR PART START		
					int flag = threadIdx.x % 4;
					/*
					if (flag == 0){
						downIndex = down;
						rightIndex = right;
					} else if (flag == 1){
						downIndex = down + step;
						rightIndex = right;
					} else if (flag == 2){
						downIndex = down;
						rightIndex = right + step;
					} else {
						downIndex = down + step;
						rightIndex = right + step;
					}
					
					rMirr = (rightIndex / step2) * step2 * 2 - rightIndex;
					rMirr = (rMirr == rightIndex) ? rightIndex : rMirr + step2;
					
					dMirr = (downIndex / step2) * step2 * 2 - downIndex;
					dMirr = (dMirr == downIndex) ? downIndex : dMirr + step2;
					
					a00ai = downIndex*N*4 + 4*rightIndex;
					a00bi = a00ai + 1;
					a00ci = a00ai + 2;
					a00di = a00ai + 3;
		
					a11ai = dMirr*N*4 + 4*rMirr;
					a11bi = a11ai + 1;
					a11ci = a11ai + 2;
					a11di = a11ai + 3;
					
					a01ai = downIndex*N*4 + 4*rMirr;
					a01bi = a01ai + 1;
					a01ci = a01ai + 2;
					a01di = a01ai + 3;
					
					a10ai = dMirr*N*4 + 4*rightIndex;
					a10bi = a10ai + 1;
					a10ci = a10ai + 2;
					a10di = a10ai + 3;
					
					a00a = a[a00ai];
					a00b = a[a00bi];
					a00c = a[a00ci];
					a00d = a[a00di];
					
					if (downIndex != dMirr && rightIndex != rMirr){
						a[a11ai] = a00a;
						a[a11bi] = -a00b;
						a[a11ci] = -a00c;
						a[a11di] = a00d;
					}
					if (downIndex != dMirr){
						a[a10ai] = a00a;
						a[a10bi] = -a00b;
						a[a10ci] = a00c;
						a[a10di] = -a00d;
					}
					if (rightIndex != rMirr){
						a[a01ai] = a00a;
						a[a01bi] = a00b;
						a[a01ci] = -a00c;
						a[a01di] = -a00d;
					}
					*/
					//MIRROR PART END
				}
			}
		}
		__syncthreads();
	}
	
}